/**
*							Instituto de Ciencias Matematicas e Computacao (ICMC)
*										Universidade de Sao Paulo (USP)
*
* Henrique Augusto Kraus Diomede
* Herik Magalhaes Serra
* Vinicius Martins dos Santos
*/


#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <curand_kernel.h>
#include <string.h>
#include <stdlib.h>

#define N 100
#define WORD_MAX 5

/**************************************************************/
/*                  Begin Function Prototypes                 */
/**************************************************************/

/*
* Faz o parser do arquivo, retorna o numero de palavras
*/
__host__ int* parser(char *buffer, int begin, int length, int *wcount);

/*
* Carregar o arquivo para memoria principal, retorna o vetor 
* caracteres
*/
__host__ char * fillBuffer(int *length,char *fileName);

/*
* Verifica se e um caracter
*/
__host__ static int isLetter(char ch);

/*
* Letra maiuscula
*/
__host__ static char toUpperCase(char ch);

/*
* Verifica se e caracter de terminacao
*/
__host__ int isTermCh(char ch);

/*
* Gera os caracteres
*/
__global__ void kernel(char *text, int len, int *wIndex, int *progress);

/*
* Verifica se o caracters esta contido no buffer
*/
__device__ int contains(char *buffer, char smallCh, char upperCh);

__constant__ char lettersDevice[26];
__constant__ int wc;

int main(void)
{
	char letters[] = {'A','B','C','D','E','F','G','H','I','J','K','L','M','N','O','P','Q','R','S','T','U','V','W','X','Y','Z',
	'a','b','c','d','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','t','u','v','w','x','y','z'}; //26 letras
	int length;
	int wordCount = 0;
	int h_count = 0;
	char *text;
	int *wIndex;
	int *dev_index;
	char *dev_text;
	int *dev_count;
	cudaEvent_t start,stop;
	float elapsedTime;

	text = fillBuffer(&length,"palavras.txt");
	wIndex = parser(text,0,length,&wordCount);

	cudaMemcpyToSymbol("lettersDevice", letters, sizeof(char) * 26 );
	cudaMemcpyToSymbol("wc", &wordCount, sizeof(int));
	
	cudaMalloc((void**)&dev_index,wordCount*sizeof(int));
	cudaMalloc((void**)&dev_text,(length+1)*sizeof(char));
	cudaMalloc((void**)&dev_count,sizeof(int));

	cudaMemcpy(dev_count,&h_count,sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_text,text,(length+1)*sizeof(char),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_index,wIndex,wordCount*sizeof(int),cudaMemcpyHostToDevice);

	
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	
	cudaEventRecord(start, 0);

	/*Kernel Launch*/
	kernel<<<64,512,(512*sizeof(int))>>>(dev_text,length,dev_index,dev_count);
	cudaThreadSynchronize();
	
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);

	cudaEventElapsedTime(&elapsedTime, start, stop);


	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	cudaMemcpy(wIndex,dev_index,wordCount*sizeof(int),cudaMemcpyDeviceToHost);	
	cudaMemcpy(&h_count,dev_count,sizeof(int),cudaMemcpyDeviceToHost);
	
	printf("Word Count %d\n",h_count);
	printf("Execution Time: %fms\n", elapsedTime);
	cudaFree(dev_index);
	cudaFree(dev_text);
	cudaFree(dev_count);
	free(text);
	scanf("%*c");
	return 0;
}


/**************************************************************/
/*                   DEVICE IMPLEMENTATIONS					  */
/**************************************************************/
__global__ void kernel(char *text, int len, int *wIndex,int *progress)
{
	extern __shared__ int blockCount[];
	char palavra[WORD_MAX+1];
	int id = threadIdx.x + blockIdx.x * blockDim.x;
	int i,j,x;
	int w,a;
	int index = id;
	char letter,upperLetter;
	int tam;
	curandState localState;
	curand_init(1234,id,0,&localState);
	blockCount[threadIdx.x] = 0;

	for(i = 0; i < N; i++)
	{

		if(index < wc-1)
		{
			tam = wIndex[index+1]-wIndex[index];
		
			for(w = 0; w < tam; w += j)
			{
				/*Copia uma palavra ou parte dela para um buffer*/
				for(j = 0; j < WORD_MAX && j < tam; j++)
				{
					palavra[j] = text[j+wIndex[index]];
				}
				palavra[j] = '\0';
			
				/*Gera caracteres aleatoriamente*/
				a = 0;
				while( a < j)
				{
					x = curand(&localState) % 26;
					letter = lettersDevice[x];
					upperLetter = lettersDevice[x+26];
					if(contains(palavra,letter,upperLetter))
						a++;
				}
			}
			/*incrementa a quantidade de palavras encontradas por
			determinada thread*/
			blockCount[threadIdx.x]++;		
		}
		/*Obtem nova posicao para indexar o vetor de palavras*/
		index += blockDim.x*gridDim.x;
	}
	__syncthreads();

	/*Atualizacao do valor global*/
	if(threadIdx.x == 0)
	{
		for(i = 1; i < blockDim.x;i++)
			blockCount[0] += blockCount[i];
		atomicAdd(progress,blockCount[0]);
	}
	__syncthreads();
}


__device__ int contains(char *buffer, char smallCh, char upperCh)
{
	int i = 0;
	while(buffer[i] != '\0')
	{
		if(buffer[i] == smallCh || buffer[i] == upperCh)
			return 1;
		else
			i++;
	}
	return 0;
}



/**************************************************************/
/*                    HOST IMPLEMENTATIONS					  */
/**************************************************************/
__host__ int* parser(char *buffer, int begin, int length, int *wcount)
{
	int lastIndex = 0;
	unsigned int end = 0;
	*wcount = 0;
	int *windex = (int*)malloc((length)*sizeof(int));

	if(buffer == NULL)
		return NULL;
	windex[0] = 0;
	(*wcount)++;

	do
	{
		while(!isLetter(buffer[begin]) && begin < length)
		{
			begin++;
			end = begin;
		}

		if(!isLetter(buffer[end]))
		{
			windex[*wcount] = lastIndex;
			(*wcount)++; 
			begin = end + 1;
		}

		if(isLetter(buffer[end]))
		{
			buffer[lastIndex] = toUpperCase(buffer[end]);
			lastIndex++;
		}
		end++;
	} while(end <= length);

	if(!isLetter(buffer[end-2]))
		(*wcount)--;

	return windex;
}

__host__ char * fillBuffer(int *length, char *fileName)
{
	FILE *fp;
	char *buffer = NULL;

	if((fp = fopen(fileName,"r")) == NULL)
		return NULL;

	fseek(fp,0,SEEK_END);
	*length = ftell(fp)/sizeof(char);
	(*length)++;
	fseek(fp,0,SEEK_SET);
	buffer = (char*)malloc((*length) * sizeof(char));
	fread(buffer,sizeof(char),--(*length),fp);
	buffer[*length] = '\0';
	fclose(fp);
	return buffer;
}

__host__ int isLetter(char ch)
{
	if(ch >= 65 && ch <= 122)
	{
		if(ch <= 90 || ch >= 97)
			return 1;
		else
			return 0;
	}
	return 0;
}

__host__ char toUpperCase(char ch)
{
	if(ch >= 97)
		return ch - 32;
	else
		return ch;
}

